/*
 *  Copyright 2008-2011 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#include <thrust/iterator/iterator_traits.h>
#include <thrust/detail/type_traits.h>
#include <thrust/functional.h>

#include <thrust/detail/backend/cuda/detail/comparison_sort.h>
#include <thrust/detail/backend/cuda/detail/numeric_sort.h>

#include <thrust/iterator/iterator_traits.h>

namespace thrust
{
namespace detail
{
namespace backend
{
namespace cuda
{
namespace dispatch
{

template<typename RandomAccessIterator,
         typename StrictWeakOrdering>
  void sort(RandomAccessIterator first,
            RandomAccessIterator last,
            StrictWeakOrdering comp,
            thrust::detail::true_type)
{
    // if comp is greater<T> then call descending sort
    typedef typename thrust::iterator_traits<RandomAccessIterator>::value_type KeyType;
    const static bool use_descending = thrust::detail::is_same<StrictWeakOrdering, typename thrust::greater<KeyType> >::value;

    if (use_descending)
    {
        thrust::detail::backend::cuda::detail::stable_descending_numeric_sort(first,last);
    }
    else
    {
        thrust::detail::backend::cuda::detail::stable_ascending_numeric_sort(first, last);
    }
}

template<typename RandomAccessIterator,
         typename StrictWeakOrdering>
  void sort(RandomAccessIterator first,
            RandomAccessIterator last,
            StrictWeakOrdering comp,
            thrust::detail::false_type)
{
    thrust::detail::backend::cuda::detail::comparison_sort(first, last, comp);
}

template<typename RandomAccessIterator,
         typename StrictWeakOrdering>
  void stable_sort(RandomAccessIterator first,
                   RandomAccessIterator last,
                   StrictWeakOrdering comp,
                   thrust::detail::true_type)
{
    // if comp is greater<T> then call descending sort
    typedef typename thrust::iterator_traits<RandomAccessIterator>::value_type KeyType;
    const static bool use_descending = thrust::detail::is_same<StrictWeakOrdering, typename thrust::greater<KeyType> >::value;

    if (use_descending)
    {
        thrust::detail::backend::cuda::detail::stable_descending_numeric_sort(first,last);
    }
    else
    {
        thrust::detail::backend::cuda::detail::stable_ascending_numeric_sort(first, last);
    }
}

template<typename RandomAccessIterator,
         typename StrictWeakOrdering>
  void stable_sort(RandomAccessIterator first,
                   RandomAccessIterator last,
                   StrictWeakOrdering comp,
                   thrust::detail::false_type)
{
    thrust::detail::backend::cuda::detail::stable_comparison_sort(first, last, comp);
}

template<typename RandomAccessIterator1,
         typename RandomAccessIterator2,
         typename StrictWeakOrdering>
  void stable_sort_by_key(RandomAccessIterator1 keys_first,
                          RandomAccessIterator1 keys_last,
                          RandomAccessIterator2 values_first,
                          StrictWeakOrdering comp,
                          thrust::detail::true_type)
{
    // if comp is greater<T> then use descending sort
    typedef typename thrust::iterator_traits<RandomAccessIterator1>::value_type KeyType;
    const static bool use_descending = thrust::detail::is_same<StrictWeakOrdering, typename thrust::greater<KeyType> >::value;

    if(use_descending)
    {
        thrust::detail::backend::cuda::detail::stable_descending_numeric_sort_by_key(keys_first, keys_last, values_first);
    }
    else
    {
        thrust::detail::backend::cuda::detail::stable_ascending_numeric_sort_by_key(keys_first, keys_last, values_first);
    }
}

template<typename RandomAccessIterator1,
         typename RandomAccessIterator2,
         typename StrictWeakOrdering>
  void stable_sort_by_key(RandomAccessIterator1 keys_first,
                          RandomAccessIterator1 keys_last,
                          RandomAccessIterator2 values_first,
                          StrictWeakOrdering comp,
                          thrust::detail::false_type)
{
    thrust::detail::backend::cuda::detail::stable_comparison_sort_by_key(keys_first, keys_last, values_first, comp);
}

} // end namespace dispatch


template<typename RandomAccessIterator,
         typename StrictWeakOrdering>
  void sort(RandomAccessIterator first,
            RandomAccessIterator last,
            StrictWeakOrdering comp)
{
    // we're attempting to launch a kernel, assert we're compiling with nvcc
    // ========================================================================
    // X Note to the user: If you've found this line due to a compiler error, X
    // X you need to compile your code using nvcc, rather than g++ or cl.exe  X
    // ========================================================================
    THRUST_STATIC_ASSERT( (depend_on_instantiation<RandomAccessIterator, THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC>::value) );

    // dispatch on whether we can use a numeric sort
    typedef typename thrust::iterator_traits<RandomAccessIterator>::value_type KeyType;
    static const bool use_numeric_sort = thrust::detail::is_arithmetic<KeyType>::value &&
                                         (thrust::detail::is_same<StrictWeakOrdering, typename thrust::less<KeyType> >::value ||
                                          thrust::detail::is_same<StrictWeakOrdering, typename thrust::greater<KeyType> >::value);


    // XXX WAR nvcc 3.0 unused variable warning
    (void) use_numeric_sort;

    thrust::detail::backend::cuda::dispatch::sort(first, last, comp,
        thrust::detail::integral_constant<bool, use_numeric_sort>());
}


template<typename RandomAccessIterator,
         typename StrictWeakOrdering>
  void stable_sort(RandomAccessIterator first,
                   RandomAccessIterator last,
                   StrictWeakOrdering comp)
{
    // we're attempting to launch a kernel, assert we're compiling with nvcc
    // ========================================================================
    // X Note to the user: If you've found this line due to a compiler error, X
    // X you need to compile your code using nvcc, rather than g++ or cl.exe  X
    // ========================================================================
    THRUST_STATIC_ASSERT( (depend_on_instantiation<RandomAccessIterator, THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC>::value) );

    // dispatch on whether we can use a numeric sort
    typedef typename thrust::iterator_traits<RandomAccessIterator>::value_type KeyType;
    static const bool use_numeric_sort = thrust::detail::is_arithmetic<KeyType>::value &&
                                         (thrust::detail::is_same<StrictWeakOrdering, typename thrust::less<KeyType> >::value ||
                                          thrust::detail::is_same<StrictWeakOrdering, typename thrust::greater<KeyType> >::value);


    // XXX WAR nvcc 3.0 unused variable warning
    (void) use_numeric_sort;

    thrust::detail::backend::cuda::dispatch::stable_sort(first, last, comp,
        thrust::detail::integral_constant<bool, use_numeric_sort>());
}

template<typename RandomAccessIterator1,
         typename RandomAccessIterator2,
         typename StrictWeakOrdering>
  void stable_sort_by_key(RandomAccessIterator1 keys_first,
                          RandomAccessIterator1 keys_last,
                          RandomAccessIterator2 values_first,
                          StrictWeakOrdering comp)
{
    // we're attempting to launch a kernel, assert we're compiling with nvcc
    // ========================================================================
    // X Note to the user: If you've found this line due to a compiler error, X
    // X you need to compile your code using nvcc, rather than g++ or cl.exe  X
    // ========================================================================
    THRUST_STATIC_ASSERT( (depend_on_instantiation<RandomAccessIterator1, THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC>::value) );

    // dispatch on whether we can use a numeric sort
    typedef typename thrust::iterator_traits<RandomAccessIterator1>::value_type KeyType;
    static const bool use_numeric_sort = thrust::detail::is_arithmetic<KeyType>::value &&
                                         (thrust::detail::is_same<StrictWeakOrdering, typename thrust::less<KeyType> >::value ||
                                          thrust::detail::is_same<StrictWeakOrdering, typename thrust::greater<KeyType> >::value);

    // XXX WAR nvcc 3.0 unused variable warning
    (void) use_numeric_sort;
    
    thrust::detail::backend::cuda::dispatch::stable_sort_by_key(keys_first, keys_last, values_first, comp,
        thrust::detail::integral_constant<bool, use_numeric_sort>());
}

} // end namespace cuda
} // end namespace backend
} // end namespace detail
} // end namespace thrust

